#include <Dynamics/Damping/velocity_damping_cuda.h>
#include <cuda_block_size_set.h>

namespace PhysLeo
{
  
template<typename T>
__global__ static void K_VelocityDampingCuda_addDampingForce(glm::tvec3<T>* vel_ptr, glm::tvec3<T>* force_ptr, T* mass_ptr, T damping_coef, int size)
{
    int pid = threadIdx.x + (blockIdx.x * blockDim.x);
    if (pid >= size) return;
    const auto vert_id = pid;

#ifndef NDEBUG
    if(pid==0)printf("add damping force before(%f,%f,%f) ",force_ptr[vert_id][0],force_ptr[vert_id][1],force_ptr[vert_id][2]);
#endif

    force_ptr[vert_id] += -damping_coef * mass_ptr[vert_id] * vel_ptr[vert_id];

#ifndef NDEBUG
    if(pid==0)printf("after(%f,%f,%f)\n",force_ptr[vert_id][0],force_ptr[vert_id][1],force_ptr[vert_id][2]);
#endif
}

template<typename T>
void VelocityDampingCuda<T>::addDampingForce(std::shared_ptr<BufferData<glm::tvec3<T>>> vel, std::shared_ptr<BufferData<glm::tvec3<T>>> force)
{
    auto mass_ptr = VelocityDampingCpu<T>::ptr_mass_->ptrGpu();
    auto vel_ptr = vel->ptrGpu();
    auto force_ptr = force->ptrGpu();

    const int vert_num = VelocityDampingCpu<T>::ptr_mass_->size();
    
    K_VelocityDampingCuda_addDampingForce<T> << <numBlocks(vert_num), BLOCK_SIZE >> > (vel_ptr, force_ptr, mass_ptr, VelocityDampingCpu<T>::velocity_damping_param_, vert_num);
    synchronCheck;
}

template class VelocityDampingCuda<float>;
template class VelocityDampingCuda<double>;

}